feat: batch KV block copies via cudaMemcpyBatchAsync in fs connector#607
Merged
Conversation
kfirtoledo
requested review from
dannyharnik,
liu-cong and
vMaroon
as code owners
github-actions
Bot
added
the
size/XL
github-actions
Bot
added
size/L
Etelis
reviewed
May 28, 2026
| // Batched DMA path: one cudaMemcpyBatchAsync covers all per-(block, layer) | ||
| // copies for the blocks in this file (num_blocks * num_tensors). | ||
| // The batch executes in stream order; ordering within the batch is unspecified. | ||
| void TensorCopier::copy_blocks_via_batch_memcpy( |
There was a problem hiding this comment.
I'm not sure it's relevant, but HIP (AMD) also supports batch memcpy out of the box, and might be worth adding as well.
|
Unsigned commits detected! Please sign your commits. For instructions on how to set up GPG/SSH signing and verify your commits, please see GitHub Documentation. |
Submit all per-(block, layer) copies in one driver call instead of N cudaMemcpyAsync calls. Enabled by default; toggle off with USE_BATCH_MEMCPY_READ / USE_BATCH_MEMCPY_WRITE=0. Requires CUDA 12.8+. Speeds up KV-cache offload writes/reads when per-layer DMA sizes are small enough that driver dispatch dominates. Signed-off-by: Kfir Toledo <kfir.toledo@ibm.com>
cudaMemcpyBatchAsync was introduced in CUDA 12.8 — guard the batch path with #if CUDA_VERSION >= 12080 and route to the per-call cudaMemcpyAsync loop below that. Default USE_BATCH_MEMCPY_* off on older toolchains so the env knob still makes sense. Also drop thread_local on the attrs/attrs_idx inputs (never mutated, no per-thread duplication needed) and move the copy_blocks dispatcher below the helpers it dispatches to. Signed-off-by: Kfir Toledo <kfir.toledo@ibm.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Replace the per-(block, layer)
cudaMemcpyAsyncloop inTensorCopierwith a singlecudaMemcpyBatchAsync(CUDA 12.8+) submission. Submits all descriptors in one driver call, removing per-call dispatch overhead.USE_BATCH_MEMCPY_READ=0/USE_BATCH_MEMCPY_WRITE=0.srcAccessOrder=ANYset oncudaMemcpyAttributes(matches vLLM'ssimple_kv_offload/cuda_mem_ops.py).#if CUDA_VERSIONhandles the failIdx out-param that CUDA 13 dropped.Measured impact (128k tokens, TP=4,
--block-size 512)Big wins on HMA models where per-layer DMAs are small; neutral on Llama/no-HMA where each per-call copy is already large enough that driver dispatch is amortized.
Test plan